﻿/*
* Nguyễn Sỹ Bách 
* Học viện kỹ thuật quân sự
* bachgist@gmail.com
*/

#include "opencv2\core\core.hpp"
#include "opencv2\highgui\highgui.hpp"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

typedef unsigned int uint;


static void help()
{
	printf("\t\tCHUONG TRINH BIEN DOI ANH KICH THUOC LON\n"
		"Doi so dong lenh nhu sau:\n"
		"[input image name] [color|gray] [threads per block(nx,ny)] [block size(w,h)] [output image name]\n"
		"vi du: D:\\flower.jpg color 16 16 32 32\n"
		"hoac: D:\\flower.jpg gray 16 32 50 100 D:\\result_flower.jpg\n");
}

#define HANDLE_ERROR(err) if (err != cudaSuccess){\
	fprintf(stderr, "%s: [%s , %d]\n", cudaGetErrorString(err),\
		__FILE__, __LINE__); exit(EXIT_FAILURE);}


#define HANDLE_NULL(ptr) if (!ptr){\
	fprintf(stderr, #ptr" = NULL: [%s , %d]\n", __FILE__,\
	__LINE__); \
	exit(EXIT_FAILURE); }

__global__ void transformationKernel(const uchar* ptr_in,
	uint height_in, uint step, uchar* ptr_out)
{
	int x, y, i;
	y = threadIdx.y + blockIdx.y * blockDim.y;
	while (y < height_in)
	{
		x = threadIdx.x + blockIdx.x * blockDim.x;
		while (x < step)
		{
			i = y * step + x;
			if (ptr_in[i] + 50 < 255)
				ptr_out[i] = ptr_in[i] + 50;
			else ptr_out[i] = 255;
			x += gridDim.x * blockDim.x;
		}
		y += gridDim.y * blockDim.y;
	}
}


void transformationWithCuda(const cv::Mat& img_in, cv::Mat& img_out,
	dim3 thread_per_block = dim3(32, 32), dim3 block_size = dim3(32,32))
{
	HANDLE_NULL(img_in.data);
	
	int channels = img_in.type() / 8 + 1;

	//Giải phóng tài nguyên, cấp phát lại bộ nhớ cho ảnh đầu ra
	if (img_out.data)
		img_out.release();
	img_out.create(img_in.rows, img_in.cols, img_in.type());
	HANDLE_NULL(img_out.data);

	cudaError_t err = cudaSuccess;
	uchar* img_in_ptr_dev = NULL;
	uchar* img_out_ptr_dev = NULL;

	// cấp phát bộ nhớ trên device
	err = cudaMalloc((void**)&img_in_ptr_dev,
		channels*img_in.rows*img_in.cols*sizeof(uchar));
	HANDLE_ERROR(err);
	err = cudaMalloc((void**)&img_out_ptr_dev,
		channels*img_out.rows*img_out.cols*sizeof(uchar));
	HANDLE_ERROR(err);

	// copy dữ liệu từ host sang device
	err = cudaMemcpy(img_in_ptr_dev, img_in.data,
		channels*img_in.rows*img_in.cols*sizeof(uchar), cudaMemcpyHostToDevice);
	HANDLE_ERROR(err);
	transformationKernel<<<thread_per_block,block_size>>>
		(img_in_ptr_dev, img_in.rows, channels*img_in.cols, img_out_ptr_dev);
	err = cudaGetLastError();
	HANDLE_ERROR(err);

	//Copy kết quả từ device sang host
	err = cudaMemcpy(img_out.data, img_out_ptr_dev,
		channels*img_out.rows*img_out.cols*sizeof(uchar), cudaMemcpyDeviceToHost);
	HANDLE_ERROR(err);

	// giải phóng tài nguyên
	cudaFree(img_in_ptr_dev);
	cudaFree(img_out_ptr_dev);
}



int main(int argc, char* argv[])
{
	// Kiểm tra số lượng đối số dòng lệnh
	if (argc < 7)
	{
		help();
		return 1;
	}

	// Đọc ảnh truyền vào
	cv::Mat img_in = cv::imread(argv[1], strcmp(argv[2], "gray")
		? CV_LOAD_IMAGE_COLOR : CV_LOAD_IMAGE_GRAYSCALE);
	HANDLE_NULL(img_in.data);
	
	printf("Anh dau vao:       %s\n", argv[1]);
	printf("Kich thuoc:        %d x %d\n", img_in.rows, img_in.cols);
	printf("Che do doc:        %s\n", argv[2]);
	
	dim3 threads_per_block(atoi(argv[3]), atoi(argv[4]));
	dim3 block_size(atoi(argv[5]), atoi(argv[6]));
	printf("threads per block: %d x %d\n",
		threads_per_block.x, threads_per_block.y);
	printf("block size:        %d x %d\n", block_size.x, block_size.y);
	

	
	cudaEvent_t start, stop;
	float elapsed_time;

	cudaEventCreate(&start);
	cudaEventRecord(start, 0);

	cv::Mat img_out;
	transformationWithCuda(img_in, img_out, threads_per_block, block_size);


	cudaEventCreate(&stop);
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsed_time, start, stop);
	printf("ElapsedTime:       %f ms\n", elapsed_time);

	if (argc > 7)
	{
		cv::imwrite(argv[7], img_out);
		printf("Da luu anh:        %s\n", argv[7]);
	}

	//Hiện ảnh lên cửa sổ
	cv::namedWindow(argv[1], cv::WINDOW_AUTOSIZE);
	cv::imshow(argv[1], img_in);

	cv::namedWindow("result image", cv::WINDOW_AUTOSIZE);
	cv::imshow("result image", img_out);
	cv::waitKey(5000);
	cv::destroyAllWindows();
	return 0;
}